//#define EQUITANT_CACHING

#include <stdlib.h>
#include <stdio.h>
#include <string.h>

#include <cuda_runtime.h>

#include "dvmgpu.h"
#include "dvmgpu_red.h"
#include "dvmgpu_copying.h"
#include "dvmgpu_distribution.h"

#ifndef _DEBUG
#define fprintf
#endif

static double get_time_internal() {
    return dvm_time();
//    struct timeval tv;
//    gettimeofday(&tv, 0);
//    return (double)tv.tv_sec + 1e-6 * tv.tv_usec;
}

// Ссылка на самый внутренний регион данных. Для подцепления к нему новых регионов
DvmGpuDataRegion *innermostRegion = 0;
int maxGridX = 65535;

static void perform_copyin_cutting(DvmGpuArrayHeaderStruct *array, Interval *cutting) {
    if (array->deviceAddr && array->havePortion) {
        perform_copy_cutting(array, cutting, cudaMemcpyHostToDevice);
    }
}

static void perform_copyout_cutting(DvmGpuArrayHeaderStruct *array, Interval *cutting) {
    if (array->deviceAddr && array->havePortion) {
        perform_copy_cutting(array, cutting, cudaMemcpyDeviceToHost);
    }
}

static void perform_copyin(DvmGpuArrayHeaderStruct *array) {
    fprintf(stderr, "perform_copyin\n");
    perform_copyin_cutting(array, 0);
}

static void perform_copyout(DvmGpuArrayHeaderStruct *array) {
    fprintf(stderr, "perform_copyout\n");
    perform_copyout_cutting(array, 0);
}

void dvmh_get_actual_whole_(long InOutDvmArray[]) {
    DvmGpuRegionRef region = innermostRegion;
    int i, newArray = 1;
    s_DISARRAY *dvmHead = (s_DISARRAY *)((SysHandle *)InOutDvmArray[0])->pP;
    DvmGpuArrayHeaderStruct *regarr;
    if (region->arrays) {
        for (i = 0; i < region->arraysCount; i++)
            if (region->arrays[i]->dvmDisArray == dvmHead) {
                newArray = 0;
                regarr = region->arrays[i];
                break;
            }
    }
    if (!newArray)
        perform_copyout(regarr);
}

void dvmh_get_actual_(long InOutDvmArray[], long InLower[], long InUpper[]) {
    //TODO: Not implemented yet
}

static void perform_copyin_slice(DvmGpuArrayHeaderStruct *array, int dim, long beginIndex, long endIndex) {
    fprintf(stderr, "perform_copyin_slice %d %ld %ld\n", dim, beginIndex, endIndex);
    if (array->deviceAddr && array->havePortion)
        perform_copy_slice(array, dim, beginIndex, endIndex, cudaMemcpyHostToDevice);
}

static void perform_copyout_slice(DvmGpuArrayHeaderStruct *array, int dim, long beginIndex, long endIndex) {
    fprintf(stderr, "perform_copyout_slice %d %ld %ld\n", dim, beginIndex, endIndex);
    if (array->deviceAddr && array->havePortion)
        perform_copy_slice(array, dim, beginIndex, endIndex, cudaMemcpyDeviceToHost);
}

static void fill_headers(DvmGpuArrayHeaderStruct *array) {
    long *header = array->header;
    header[0] = (long)array;
    if (array->deviceAddr && array->havePortion) {
        int i;
        long collector = 1;
        long offset = 0;
        for (i = array->rank - 1; i >= 0; i--) {
            offset += collector * array->havePortion[i][0];
            header[i + 1] = collector;
            collector *= array->havePortion[i][1] - array->havePortion[i][0] + 1;
        }
        header[array->rank] = ((long)array->deviceAddr - header[array->rank + 2]) / array->typeSize;
        header[array->rank + 1] = ((long)array->deviceAddr - header[array->rank + 2]) / array->typeSize - offset;
    } else {
        int i;
        for (i = 0; i < array->rank + 1; i++) {
            header[1 + i] = 0;
        }
    }
    int i;
    for (i = 0; i < array->additionalHeadersCount; i++) {
        long *header2 = array->additionalHeaders[i];
        int j;
        for (j = 0; j < array->rank + 1; j++)
            header2[j] = header[j];
        header2[array->rank + 1] = (header[array->rank + 2] - header2[array->rank + 2]) / array->typeSize + header[array->rank + 1];
    }
}

void dvmcommrank_(int *);
// Делает всю необходимую инициализацию. Должна быть вызвана до регистрации первого региона(данных или вычислений не важно).
long initdvm_gpu_() {
    int rank;
    int cudaDevice;
    dvmcommrank_(&rank);
    cudaDevice = rank % 3;
    cudaSetDevice(cudaDevice);
    innermostRegion = 0;
// TODO: request device memory amount, shared memory amount. Restrictions for grid size and block size
    fprintf(stderr, "OK1 rank=%d\n", rank);
}

static DvmGpuRegionRef crt_region(enum RegionType regionType) {
    DvmGpuRegionRef newRegion;
    newRegion = (DvmGpuRegionRef)malloc(sizeof(DvmGpuRegion));
    newRegion->regionType = regionType;
    newRegion->parentRegion = innermostRegion;
    if (newRegion->parentRegion && newRegion->parentRegion->arraysCount) {
        newRegion->arraysCount = newRegion->parentRegion->arraysCount;
        newRegion->arrays = (DvmGpuArrayHeaderStruct **)malloc(sizeof(DvmGpuArrayHeaderStruct *) * newRegion->arraysCount);
        memcpy(newRegion->arrays, newRegion->parentRegion->arrays, newRegion->arraysCount * sizeof(DvmGpuArrayHeaderStruct *));
    } else {
        newRegion->arrays = 0;
        newRegion->arraysCount = 0;
    }
    newRegion->myArrays = 0;
    newRegion->myArraysCount = 0;
    newRegion->additionalHeadersCount = 0;
    newRegion->additionalHeaders = 0;
    innermostRegion = newRegion;
    return newRegion;
}

// Регистрирует вычислительный регион
long crt_region_gpu_() {
    return (long)crt_region(RT_COMPUTE);
}

// Регистрирует регион данных
long crt_data_region_gpu_() {
    return (long)crt_region(RT_DATA);
}

// Регистрирует DVM-массив в регионе, заполняет начальными данными заголовок GPU-массива (заголовок лежит на хосте)
// При вызове для региона данных - выполняет аллокацию и, если нужно, копирование
// При вызове для региона вычислений только регистрирует. Аллокации будут в startpl_gpu_ и loop_gpu_
void crtda_gpu_(DvmGpuRegionRef *InRegionGpu, long InDvmArray[], long OutDvmGpuArray[], void *InDeviceBaseAddr, long *InCopyinFlag, long *InCopyoutFlag) {
    DvmGpuRegionRef region = *InRegionGpu;
    int i, newArray = 1;
    s_DISARRAY *dvmHead = (s_DISARRAY *)((SysHandle *)InDvmArray[0])->pP;
    DvmGpuArrayHeaderStruct *regarr;
    if (region->arrays) {
        for (i = 0; i < region->arraysCount; i++)
            if (region->arrays[i]->dvmDisArray == dvmHead) {
                newArray = 0;
                regarr = region->arrays[i];
                break;
            }
    }

    if (newArray) {
// Register array in region
        region->myArraysCount++;
        region->arraysCount++;
        if (!region->arrays) {
            region->myArrays = (DvmGpuArrayHeaderStruct **)malloc(region->myArraysCount * sizeof(DvmGpuArrayHeaderStruct *));
            region->arrays = (DvmGpuArrayHeaderStruct **)malloc(region->arraysCount * sizeof(DvmGpuArrayHeaderStruct *));
        } else if (!region->myArrays) {
            region->myArrays = (DvmGpuArrayHeaderStruct **)malloc(region->myArraysCount * sizeof(DvmGpuArrayHeaderStruct *));
        } else {
            region->myArrays = (DvmGpuArrayHeaderStruct **)realloc(region->myArrays, region->myArraysCount * sizeof(DvmGpuArrayHeaderStruct **));
            region->arrays = (DvmGpuArrayHeaderStruct **)realloc(region->arrays, region->arraysCount * sizeof(DvmGpuArrayHeaderStruct **));
        }
        regarr = (DvmGpuArrayHeaderStruct *)malloc(sizeof(DvmGpuArrayHeaderStruct));
        region->myArrays[region->myArraysCount - 1] = regarr;
        region->arrays[region->arraysCount - 1] = regarr;
// Fill header
        regarr->header = OutDvmGpuArray;
        regarr->copyinFlag = *InCopyinFlag;
        regarr->copyoutFlag = *InCopyoutFlag;
        regarr->deviceAddr = 0;
        regarr->additionalHeaders = 0;
        regarr->additionalHeadersCount = 0;

        regarr->dvmDisArray = dvmHead;
        regarr->typeSize = dvmHead->TLen;
        regarr->rank = dvmHead->Space.Rank;//ArrBlock.Block.Rank;
        OutDvmGpuArray[regarr->rank + 2] = (long)InDeviceBaseAddr;
        regarr->havePortion = 0;
        regarr->needPortion = 0;
        if (region->regionType == RT_DATA || region->regionType == RT_COMPUTE) {
// Allocate all data at entrance
            long *dvmheader = (long *)dvmHead->HandlePtr->HeaderPtr;
            long elemCount = 1;
            regarr->havePortion = (Interval *)malloc(sizeof(Interval) * regarr->rank);
            for (i = 0; i < regarr->rank; i++) {
                regarr->havePortion[i][0] = dvmHead->ArrBlock.Block.Set[i].Lower + dvmheader[regarr->rank + i + 2];
                regarr->havePortion[i][1] = dvmHead->ArrBlock.Block.Set[i].Upper + dvmheader[regarr->rank + i + 2];
                fprintf(stderr, "[%ld..%ld]", regarr->havePortion[i][0], regarr->havePortion[i][1]);
                elemCount *= regarr->havePortion[i][1] - regarr->havePortion[i][0] + 1;
            }
            cudaMalloc(&regarr->deviceAddr, elemCount * regarr->typeSize);
            fprintf(stderr, " allocating data for %d elements\n", elemCount);
            if (regarr->copyinFlag) {
                fill_headers(regarr);
                perform_copyin(regarr);
            }
        }
    } else {
        int isNewHeader = regarr->header != OutDvmGpuArray;
        for (i = 0; isNewHeader && i < regarr->additionalHeadersCount; i++)
            if (regarr->additionalHeaders[i] == OutDvmGpuArray)
                isNewHeader = 0;
        if (isNewHeader) {
            regarr->additionalHeadersCount++;
            if (!regarr->additionalHeaders)
                regarr->additionalHeaders = (long **)malloc(sizeof(long *) * regarr->additionalHeadersCount);
            else
                regarr->additionalHeaders = (long **)realloc(regarr->additionalHeaders, sizeof(long *) * regarr->additionalHeadersCount);
            regarr->additionalHeaders[regarr->additionalHeadersCount - 1] = OutDvmGpuArray;
            region->additionalHeadersCount++;
            if (!region->additionalHeaders)
                region->additionalHeaders = (long **)malloc(sizeof(long *) * region->additionalHeadersCount);
            else
                region->additionalHeaders = (long **)realloc(region->additionalHeaders, sizeof(long *) * region->additionalHeadersCount);
            region->additionalHeaders[region->additionalHeadersCount - 1] = OutDvmGpuArray;
        }
    }
    OutDvmGpuArray[regarr->rank + 2] = (long)InDeviceBaseAddr;
    fill_headers(regarr);
    if (newArray)
        fprintf(stderr, "OK3 array=%ld rank=%d typesize=%d copyin=%d copyout=%d\n", (long)regarr, regarr->rank, regarr->typeSize, regarr->copyinFlag,
                regarr->copyoutFlag);
}

long strtsh_gpu_(DvmGpuComputeRegionRef *InComputeRegionGpu, ShadowGroupRef *ShadowGroupRefPtr) {
    s_BOUNDGROUP *bg = (s_BOUNDGROUP *)((SysHandle *)*ShadowGroupRefPtr)->pP;
    DvmGpuComputeRegion *region = *InComputeRegionGpu;
    DvmGpuShadow *shadow = (DvmGpuShadow *)malloc(sizeof(DvmGpuShadow));
    shadow->slicesCount = 0;
    int i, totalMax = 0;;
    for (i = 0; i < bg->NewArrayColl.Count; i++)
        totalMax += ((s_DISARRAY *)bg->NewArrayColl.List[i])->Space.Rank * 2;
    shadow->sliceDescs = (SliceDesc *)malloc(totalMax * sizeof(SliceDesc));
    for (i = 0; i < bg->NewArrayColl.Count; i++) {
        s_DISARRAY *dvmHead = bg->NewArrayColl.List[i];
        int haveArray = 0;
        DvmGpuArrayHeaderStruct *array;
        if (region->arrays) {
            int j;
            for (j = 0; j < region->arraysCount; j++)
                if (region->arrays[j]->dvmDisArray == dvmHead) {
                    haveArray = 1;
                    array = region->arrays[j];
                    break;
                }
        }
        if (haveArray) {
            s_SHDWIDTH *shdWidth = bg->NewShdWidthColl.List[i];
            int j;
            for (j = 0; j < array->rank; j++) {
                long starti = ((long *)dvmHead->HandlePtr->HeaderPtr)[array->rank + j + 2];
                long beginIndex, endIndex;

                beginIndex = dvmHead->Block.Set[j].Lower;
                endIndex = dvmHead->Block.Set[j].Lower + shdWidth->ResHighShdWidth[j] - 1;
                if (endIndex > dvmHead->ArrBlock.Block.Set[j].Upper)
                    endIndex = dvmHead->ArrBlock.Block.Set[j].Upper;
                if (endIndex >= beginIndex && beginIndex > 0) {
                    perform_copyout_slice(array, j + 1, beginIndex + starti, endIndex + starti);
                }

                beginIndex = dvmHead->Block.Set[j].Upper - shdWidth->ResLowShdWidth[j] + 1;
                endIndex = dvmHead->Block.Set[j].Upper;
                if (beginIndex < dvmHead->ArrBlock.Block.Set[j].Lower)
                    beginIndex = dvmHead->ArrBlock.Block.Set[j].Lower;
                if (endIndex >= beginIndex && endIndex < dvmHead->Space.Size[j] - 1) {
                    perform_copyout_slice(array, j + 1, beginIndex + starti, endIndex + starti);
                }

                beginIndex = dvmHead->Block.Set[j].Lower - shdWidth->ResLowShdWidth[j];
                endIndex = dvmHead->Block.Set[j].Lower - 1;
                if (beginIndex < dvmHead->ArrBlock.Block.Set[j].Lower)
                    beginIndex = dvmHead->ArrBlock.Block.Set[j].Lower;
                if (endIndex >= beginIndex && beginIndex >= 0) {
                    shadow->sliceDescs[shadow->slicesCount].array = array;
                    shadow->sliceDescs[shadow->slicesCount].dim = j + 1;
                    shadow->sliceDescs[shadow->slicesCount].beginIndex = beginIndex + starti;
                    shadow->sliceDescs[shadow->slicesCount].endIndex = endIndex + starti;
                    shadow->slicesCount++;
                }

                beginIndex = dvmHead->Block.Set[j].Upper + 1;
                endIndex = dvmHead->Block.Set[j].Upper + shdWidth->ResHighShdWidth[j];
                if (endIndex > dvmHead->ArrBlock.Block.Set[j].Upper)
                    endIndex = dvmHead->ArrBlock.Block.Set[j].Upper;
                if (endIndex >= beginIndex && endIndex < dvmHead->Space.Size[j]) {
                    shadow->sliceDescs[shadow->slicesCount].array = array;
                    shadow->sliceDescs[shadow->slicesCount].dim = j + 1;
                    shadow->sliceDescs[shadow->slicesCount].beginIndex = beginIndex + starti;
                    shadow->sliceDescs[shadow->slicesCount].endIndex = endIndex + starti;
                    shadow->slicesCount++;
                }
            }
        }
    }
    return (long)shadow;
}

void donesh_gpu_(DvmGpuShadowRef *InDvmGpuShadowRefPtr) {
    DvmGpuShadow *shadow = *InDvmGpuShadowRefPtr;
    int i;
    for (i = 0; i < shadow->slicesCount; i++)
        perform_copyin_slice(shadow->sliceDescs[i].array, shadow->sliceDescs[i].dim, shadow->sliceDescs[i].beginIndex, shadow->sliceDescs[i].endIndex);
    safe_free((void **)&shadow->sliceDescs);
    safe_free((void **)&shadow);
}

// Создание параллельного DVM_GPU цикла
long crtpl_gpu_(DvmGpuComputeRegionRef *InComputeRegionGpu, LoopRef *InDvmLoop, long *InFirstInRegionFlag, long *InLastInRegionFlag) {
    DvmGpuLoopRef newLoop = (DvmGpuLoopRef)malloc(sizeof(DvmGpuLoop));
    newLoop->region = *InComputeRegionGpu;
    newLoop->firstLoopInRegion = *InFirstInRegionFlag;
    newLoop->lastLoopInRegion = *InLastInRegionFlag;

    s_PARLOOP *dvmLoop = (s_PARLOOP *)((SysHandle *)(*InDvmLoop))->pP;
    newLoop->dvmLoop = dvmLoop;
    newLoop->dimension = dvmLoop->Rank;
    newLoop->loopBounds = (LoopBounds *)malloc(sizeof(LoopBounds) * newLoop->dimension);
    int i;
    for (i = 0; i < newLoop->dimension; i++) {
        newLoop->loopBounds[i][0] = (*dvmLoop->MapList[i].InitIndexPtr);
        newLoop->loopBounds[i][1] = (*dvmLoop->MapList[i].LastIndexPtr);
        newLoop->loopBounds[i][2] = (*dvmLoop->MapList[i].StepPtr);
    }
    newLoop->counter = -1;
    newLoop->reductions = 0;
    newLoop->reductionsCount = 0;

    for (i = 0; i < newLoop->region->arraysCount; i++)
        safe_free((void **)&newLoop->region->arrays[i]->needPortion);

    return (long)newLoop;
}

// InAxisArray		массив i-й элемент которого содержит номер измерения параллельного цикла, соответствующего (i+1)-му измерению распределённого массива.
// InCoeffArray		массив, i-й элемент которого содержит коэффициент индексной переменной линейного правила выборки для (i+1)-го измерения
//   распределённого массива.
// InConstArray		массив, i-й элемент которого содержит константу линейного правила  выборки для (i+1)-го измерения распределённого массива.
// Добавление ссылки на массив.
void add_array_reference_gpu_(DvmGpuLoopRef *InDvmGpuLoop, long InDvmGpuArray[], long InAxisArray[], long InCoeffArray[], long InConstArray[]) {
    DvmGpuLoopRef loop = (DvmGpuLoopRef)((long *)InDvmGpuLoop)[0];
    DvmGpuArrayHeaderStruct *array = (DvmGpuArrayHeaderStruct *)InDvmGpuArray[0];
    int wasEmptyFlag = 0;
    if (!array->needPortion) {
        array->needPortion = (Interval *)malloc(sizeof(Interval) * array->rank);
        wasEmptyFlag = 1;
    }
    int j;
    for (j = 0; j < 2; j++) {
        int i;
        for (i = 0; i < array->rank; i++) {
            long ind = loop->loopBounds[InAxisArray[i] - 1][j] * InCoeffArray[i] + InConstArray[i] +
                    ((long *)array->dvmDisArray->HandlePtr->HeaderPtr)[array->rank + i + 2];
            if (wasEmptyFlag || ind < array->needPortion[i][0]) {
                array->needPortion[i][0] = ind;
            }
            if (wasEmptyFlag || ind > array->needPortion[i][1]) {
                array->needPortion[i][1] = ind;
            }
        }
        wasEmptyFlag = 0;
    }
    for (j = 0; j < array->rank; j++)
        fprintf(stderr, "needPortion[%d] = [%ld, %ld]\n", j, array->needPortion[j][0], array->needPortion[j][1]);
}

// Опрос возможности выполнения на GPU цикла. Если возможно, то иницииализируются данные для управления ходом вычисления на GPU цикла и возвращается 1.
// Если невозможно, возвращает нуль и в случае, если есть на ГПУ данные, не выкачанные из соображений, что пока рано, то их надо выкачать сейчас,
// чтобы обеспечить счёт на ЦПУ всеми сделанными изменениями.
long startpl_gpu_(DvmGpuLoopRef *InDvmGpuLoop) {
    DvmGpuLoopRef loop = (DvmGpuLoopRef)((long *)InDvmGpuLoop)[0];

    loop->counter = 0;
    return 1;

    DvmGpuComputeRegionRef region = loop->region;
    long totalNeed = 0;
    long totalHave = 0;
    long inscribNeed = 0;
    long inscribHave = 0;
    long noninsHave = 0;

    int i;
    for (i = 0; i < region->arraysCount; i++) {
        DvmGpuArrayHeaderStruct *array = region->arrays[i];
        if (array->needPortion || array->havePortion) {
            long needElemCount = !!array->needPortion;
            long haveElemCount = !!array->havePortion;
            int j, inscribFlag = 1;
            for (j = 0; j < array->rank; j++) {
                if (array->needPortion)
                    needElemCount *= array->needPortion[j][1] - array->needPortion[j][0] + 1;
                if (array->havePortion) {
                    haveElemCount *= array->havePortion[j][1] - array->havePortion[j][0] + 1;
                    if (array->needPortion && (array->needPortion[j][0] < array->havePortion[j][0] || array->needPortion[j][1] > array->havePortion[j][1]))
                        inscribFlag = 0;
                }
            }
            totalNeed += needElemCount * array->typeSize;
            totalHave += haveElemCount * array->typeSize;
            if (array->havePortion && inscribFlag) {
                inscribNeed += needElemCount * array->typeSize;
                inscribHave += haveElemCount * array->typeSize;
            } else if (array->needPortion && array->havePortion) {
                noninsHave += haveElemCount * array->typeSize;
            }
        }
    }
    size_t freeMemSize, totalMemSize;
    cudaMemGetInfo(&freeMemSize, &totalMemSize);
    fprintf(stderr, "totalNeed=%ld totalHave=%ld inscribNeed=%ld noinsHave=%ld free=%ld/%ld\n", totalNeed, totalHave, inscribNeed, noninsHave, freeMemSize,
            totalMemSize);
    if (freeMemSize + noninsHave >= totalNeed - inscribNeed) {
        for (i = 0; i < region->arraysCount; i++) {
            DvmGpuArrayHeaderStruct *array = region->arrays[i];
            if (array->havePortion && array->needPortion) {
                int j, inscribFlag = 1;
                long needElemCount = 1;
                for (j = 0; j < array->rank; j++) {
                    needElemCount *= array->needPortion[j][1] - array->needPortion[j][0] + 1;
                    if (array->needPortion[j][0] < array->havePortion[j][0] || array->needPortion[j][1] > array->havePortion[j][1])
                        inscribFlag = 0;
                }
                if (!inscribFlag) {
                    if (array->deviceAddr) {
                        int isMyArray = 0;
                        for (j = 0; j < region->myArraysCount; j++)
                            if (region->myArrays[j] == array) {
                                isMyArray = 1;
                                break;
                            }
                        if (!isMyArray || array->copyoutFlag || !loop->lastLoopInRegion)
                            perform_copyout(array);
                        cudaFree(array->deviceAddr);
                        array->deviceAddr = 0;
                    }
                    safe_free((void **)&array->havePortion);
                }
            }
        }
        for (i = 0; i < region->arraysCount; i++) {
            DvmGpuArrayHeaderStruct *array = region->arrays[i];
            if (!array->havePortion && array->needPortion) {
                int j;
                long needElemCount = 1;
                for (j = 0; j < array->rank; j++) {
                    needElemCount *= array->needPortion[j][1] - array->needPortion[j][0] + 1;
                }
                if (array->deviceAddr)
                    cudaFree(array->deviceAddr);
                cudaMalloc(&array->deviceAddr, needElemCount * array->typeSize);
                fprintf(stderr, "devad=%ld elc=%ld\n", array->deviceAddr, needElemCount);
                array->havePortion = (Interval *)malloc(sizeof(Interval) * array->rank);
                for (j = 0; j < array->rank; j++) {
                    int k;
                    for (k = 0; k < 2; k++)
                        array->havePortion[j][k] = array->needPortion[j][k];
                }
                fill_headers(array);
                int isMyArray = 0;
                for (j = 0; j < region->myArraysCount; j++)
                    if (region->myArrays[j] == array) {
                        isMyArray = 1;
                        break;
                    }
                if (!isMyArray || array->copyinFlag || !loop->firstLoopInRegion)
                    perform_copyin(array);
            }
        }
        loop->counter = 0;
        fprintf(stderr, "OK6\n");
        return 1;
    }
    fprintf(stderr, "neOK6\n");
    return 0;
}

static void finish_pl(DvmGpuLoopRef loop) {
    if (loop)
        safe_free((void **)&loop->loopBounds);
    if (loop->reductions) {
        int i;
        for (i = 0; i < loop->reductionsCount; i++) {
            if (loop->reductions[i]->gpuMem)
                cudaFree(loop->reductions[i]->gpuMem);
            if (loop->reductions[i]->gpuLocMem)
                cudaFree(loop->reductions[i]->gpuLocMem);
            safe_free((void **)&loop->reductions[i]);
        }
        safe_free((void **)&loop->reductions);
    }
    safe_free((void **)&loop);
}

// Функция для нормализации состояния после того, как startpl_gpu_ вернуль нуль и цикл исполнился средствами хоста
void end_host_exec_gpu_(DvmGpuLoopRef *InDvmGpuLoop) {
    finish_pl(*InDvmGpuLoop);
    fprintf(stderr, "something strange\n");
}

// Регистрация редукции в цикле
void insred_gpu_(DvmGpuLoopRef *InDvmGpuLoop, RedRef *InRedRefPtr, void *InDeviceArrayBaseAddr, void *InDeviceLocBaseAddr, AddrType *ArrayOffsetPtr, AddrType *LocOffsetPtr) {
    DvmGpuLoopRef loop = *InDvmGpuLoop;
    if (!loop->reductions) {
        loop->reductionsCount = 1;
        loop->reductions = (Reduction **)malloc(sizeof(Reduction *));
    } else {
        loop->reductionsCount++;
        loop->reductions = (Reduction **)realloc(loop->reductions, sizeof(Reduction *) * loop->reductionsCount);
    }
    Reduction *reduction = (Reduction *)malloc(sizeof(Reduction));
    loop->reductions[loop->reductionsCount - 1] = reduction;
    s_REDVAR *dvmRVar = (s_REDVAR *)((SysHandle *)(*InRedRefPtr))->pP;
    reduction->dvmRVar = dvmRVar;
    reduction->array = dvmRVar->Mem;
    reduction->arrayLength = dvmRVar->VLength;
    reduction->arrayElementType = dvmRVar->VType;
    reduction->arrayElementSize = dvmRVar->RedElmLength;
    reduction->arrayBackup = 0;
    reduction->locArray = dvmRVar->LocMem;
    reduction->locElementSize = dvmRVar->LocElmLength;
    reduction->funcNumber = dvmRVar->Func == rf_MAXLOC ? rf_MAX : (dvmRVar->Func == rf_MINLOC ? rf_MIN : dvmRVar->Func);
    reduction->gpuMem = 0;
    reduction->gpuLocMem = 0;
    reduction->arrayBaseAddr = InDeviceArrayBaseAddr;
    reduction->locBaseAddr = InDeviceLocBaseAddr;
    reduction->gpuOffsetPtr = (IndexType *)*ArrayOffsetPtr;
    reduction->gpuLocOffsetPtr = (IndexType *)*LocOffsetPtr;
}

#define CALL_REDFUNC(func, type) { red_##func##_##type(items, (type *)reduction->gpuMem, reduction->arrayLength, reduction->gpuLocMem, \
 reduction->locElementSize, (type *)reduction->array, reduction->locArray); catched = 1; }
static void finish_reduction(int items, Reduction *reduction) {
    if (reduction->arrayBackup) {
        memcpy(reduction->array, reduction->arrayBackup, reduction->arrayLength * reduction->arrayElementSize);
        safe_free((void **)&reduction->arrayBackup);
    }
    fprintf(stderr, "items=%d\n", items);
    int catched = 0;
//    if (0)
    switch(reduction->arrayElementType) {
    case rt_CHAR:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, char); break;
        case rf_MULT: CALL_REDFUNC(mult, char); break;
        case rf_MAX: CALL_REDFUNC(max, char); break;
        case rf_MIN: CALL_REDFUNC(min, char); break;
        case rf_AND: CALL_REDFUNC(and, char); break;
        case rf_OR: CALL_REDFUNC(or, char); break;
        case rf_XOR: CALL_REDFUNC(xor, char); break;
        case rf_EQU: CALL_REDFUNC(equ, char); break;
        case rf_NE: CALL_REDFUNC(ne, char); break;
        case rf_EQ: CALL_REDFUNC(eq, char); break;
        }
        break;
    case rt_INT:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, int); break;
        case rf_MULT: CALL_REDFUNC(mult, int); break;
        case rf_MAX: CALL_REDFUNC(max, int); break;
        case rf_MIN: CALL_REDFUNC(min, int); break;
        case rf_AND: CALL_REDFUNC(and, int); break;
        case rf_OR: CALL_REDFUNC(or, int); break;
        case rf_XOR: CALL_REDFUNC(xor, int); break;
        case rf_EQU: CALL_REDFUNC(equ, int); break;
        case rf_NE: CALL_REDFUNC(ne, int); break;
        case rf_EQ: CALL_REDFUNC(eq, int); break;
        }
        break;
    case rt_LONG:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, long); break;
        case rf_MULT: CALL_REDFUNC(mult, long); break;
        case rf_MAX: CALL_REDFUNC(max, long); break;
        case rf_MIN: CALL_REDFUNC(min, long); break;
        case rf_AND: CALL_REDFUNC(and, long); break;
        case rf_OR: CALL_REDFUNC(or, long); break;
        case rf_XOR: CALL_REDFUNC(xor, long); break;
        case rf_EQU: CALL_REDFUNC(equ, long); break;
        case rf_NE: CALL_REDFUNC(ne, long); break;
        case rf_EQ: CALL_REDFUNC(eq, long); break;
        }
        break;
    case rt_FLOAT:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, float); break;
        case rf_MULT: CALL_REDFUNC(mult, float); break;
        case rf_MAX: CALL_REDFUNC(max, float); break;
        case rf_MIN: CALL_REDFUNC(min, float); break;
        }
        break;
    case rt_DOUBLE:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, double); break;
        case rf_MULT: CALL_REDFUNC(mult, double); break;
        case rf_MAX: CALL_REDFUNC(max, double); break;
        case rf_MIN: CALL_REDFUNC(min, double); break;
        }
        break;
    case rt_FLOAT_COMPLEX:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, float_complex); break;
        case rf_MULT: CALL_REDFUNC(mult, float_complex); break;
        }
        break;
    case rt_DOUBLE_COMPLEX:
        switch (reduction->funcNumber) {
        case rf_SUM: CALL_REDFUNC(sum, double_complex); break;
        case rf_MULT: CALL_REDFUNC(mult, double_complex); break;
        }
        break;
    }
    if (reduction->locArray && reduction->locElementSize == 8)
        fprintf(stderr, "%f %d %d\n", *(float *)reduction->array, *(int *)reduction->locArray, *((int *)reduction->locArray + 1));
    if (!catched)
        fprintf(stderr, "something strange. not found reduction function to call\n");
}
#undef CALL_REDFUNC

// Запрос на подготовку информации для запуска ядра для очередного кванта. Возвращает не ноль, если требуется выполнения ядра и ноль, если обработка окончена.
// Определяет: порцию dvm-цикла для обработки, порции всех зарегистрированных в регионе массивов, решетку блоков, размеры блоков, порции dvm-цикла для
// каждого блока. Также выполняет все необходимые аллокации DVM_GPU массивов и переписи данных туда и обратно(если необходимо). Последним своим
// вызовом(который возвращает нуль) завершает GPU-цикл, освобождает структуры управления этим циклом.
/*long loop_gpu_(DvmGpuLoopRef *InDvmGpuLoop, dim3 *OutBlocks, dim3 *OutThreads, void *InDeviceBaseAddr, IndexType *InOutBlocksOffs) {
    DvmGpuLoopRef loop = *InDvmGpuLoop;

    if (loop->counter == 0) {
        int resBlocks;
        IndexType *deviceBlocksInfo;
        get_distribution(loop->dimension, loop->loopBounds, OutThreads, &resBlocks, &deviceBlocksInfo);
        *InOutBlocksOffs = ((long)deviceBlocksInfo - (long)InDeviceBaseAddr) / sizeof(IndexType) - 1;
        if (loop->reductions) {
            int i;
            for (i = 0; i < loop->reductionsCount; i++) {
                cudaMalloc((void **)&loop->reductions[i]->gpuMem,
                        loop->reductions[i]->arrayLength * loop->reductions[i]->arrayElementSize * resBlocks);
                *loop->reductions[i]->gpuOffsetPtr = ((long)loop->reductions[i]->gpuMem - (long)loop->reductions[i]->arrayBaseAddr) /
                        loop->reductions[i]->arrayElementSize - 1;
                if (loop->reductions[i]->gpuLocOffsetPtr) {
                    cudaMalloc((void **)&loop->reductions[i]->gpuLocMem,
                            loop->reductions[i]->arrayLength * loop->reductions[i]->locElementSize * resBlocks);
                    *loop->reductions[i]->gpuLocOffsetPtr = ((long)loop->reductions[i]->gpuLocMem - (long)loop->reductions[i]->locBaseAddr) /
                            loop->reductions[i]->locElementSize - 1;
                }
                if (loop->reductions[i]->funcNumber == rf_SUM || loop->reductions[i]->funcNumber == rf_MULT || loop->reductions[i]->funcNumber == rf_XOR) {
                    loop->reductions[i]->arrayBackup = (char *)malloc(loop->reductions[i]->arrayLength * loop->reductions[i]->arrayElementSize);
                    memcpy(loop->reductions[i]->arrayBackup, loop->reductions[i]->array,
                            loop->reductions[i]->arrayLength * loop->reductions[i]->arrayElementSize);
                    if (loop->reductions[i]->funcNumber == rf_MULT) {
                        int j;
                        for (j = 0; j < loop->reductions[i]->arrayLength; j++) {
                            char *ptr = loop->reductions[i]->array + j * loop->reductions[i]->arrayElementSize;
                            switch (loop->reductions[i]->arrayElementType) {
                            case rt_CHAR: *(char *)ptr = 1; break;
                            case rt_INT: *(int *)ptr = 1; break;
                            case rt_LONG: *(long *)ptr = 1; break;
                            case rt_FLOAT: *(float *)ptr = 1.0f; break;
                            case rt_DOUBLE: *(double *)ptr = 1.0; break;
                            case rt_FLOAT_COMPLEX: *(float *)ptr = 1.0f; *((float *)ptr + 1) = 0.0f; break;
                            case rt_DOUBLE_COMPLEX: *(double *)ptr = 1.0; *((double *)ptr + 1) = 0.0; break;
                            }
                        }
                    } else {
                        int j;
                        for (j = 0; j < loop->reductions[i]->arrayLength; j++) {
                            char *ptr = loop->reductions[i]->array + j * loop->reductions[i]->arrayElementSize;
                            switch (loop->reductions[i]->arrayElementType) {
                            case rt_CHAR: *(char *)ptr = 0; break;
                            case rt_INT: *(int *)ptr = 0; break;
                            case rt_LONG: *(long *)ptr = 0; break;
                            case rt_FLOAT: *(float *)ptr = 0.0f; break;
                            case rt_DOUBLE: *(double *)ptr = 0.0; break;
                            case rt_FLOAT_COMPLEX: *(float *)ptr = 0.0f; *((float *)ptr + 1) = 0.0f; break;
                            case rt_DOUBLE_COMPLEX: *(double *)ptr = 0.0; *((double *)ptr + 1) = 0.0; break;
                            }
                        }
                    }
                }
            }
        }
        loop->overallBlocks = resBlocks;
        loop->restBlocks = resBlocks;
        loop->latestBlocks = 0;
        loop->counter++;
    }
    if (loop->counter == 1) {
        if (loop->restBlocks <= 0) {
            loop->counter++;
        } else {
            int maxBlocks = maxGridX;
            int toExec = loop->restBlocks <= maxBlocks ? loop->restBlocks : (loop->restBlocks / 2 <= maxBlocks ? loop->restBlocks / 2 : maxBlocks);
            *InOutBlocksOffs += loop->latestBlocks * loop->dimension * 3;
            if (loop->reductions) {
                int i;
                for (i = 0; i < loop->reductionsCount; i++) {
                    *loop->reductions[i]->gpuOffsetPtr += loop->latestBlocks * loop->reductions[i]->arrayLength;
                    if (loop->reductions[i]->gpuLocOffsetPtr)
                        *loop->reductions[i]->gpuLocOffsetPtr += loop->latestBlocks * loop->reductions[i]->arrayLength;
                }
            }
            OutBlocks->x = toExec;
            OutBlocks->y = 1;
            OutBlocks->z = 1;
            loop->latestBlocks = toExec;
            loop->restBlocks -= toExec;
            fprintf(stderr, "OK7 block=(%d,%d,%d) grid=(%d,%d,%d)\n", OutThreads->x, OutThreads->y, OutThreads->z, OutBlocks->x, OutBlocks->y, OutBlocks->z);
            return 1;
        }
    }
    if (loop->counter == 2) {
        *InOutBlocksOffs = 0;
        if (loop->reductions) {
            int i;
            for (i = 0; i < loop->reductionsCount; i++)
                finish_reduction(loop->overallBlocks, loop->reductions[i]);
        }
        finish_pl(*InDvmGpuLoop);
        fprintf(stderr, "OK8\n");
        return 0;
    }
    fprintf(stderr, "neOK78\n");
    return 0;
}

long loop_ns_gpu_(DvmGpuLoopRef *InDvmGpuLoop, dim3 *OutBlocks, dim3 *OutThreads, long OutLowValues[], long OutHighValues[]) {
    DvmGpuLoopRef loop = *InDvmGpuLoop;

    if (loop->counter == 0) {
        int turns, threads;
        get_distribution_ns(OutThreads);
        threads = OutThreads->x * OutThreads->y * OutThreads->z;
        turns = 1;
        int i;
        for (i = 0; i < loop->dimension; i++) {
            turns *= (loop->loopBounds[i][1] - loop->loopBounds[i][0]) / loop->loopBounds[i][2] + 1;
            OutLowValues[i] = OutHighValues[i] = loop->loopBounds[i][0];
        }
        OutHighValues[loop->dimension - 1] -= loop->loopBounds[loop->dimension - 1][2];
        loop->restBlocks = (turns + threads - 1) / threads;
        loop->latestBlocks = 0;
        loop->counter++;
    }
    if (loop->counter == 1) {
        if (loop->restBlocks <= 0) {
            loop->counter++;
        } else {
            int maxBlocks = maxGridX;
            int threads = OutThreads->x * OutThreads->y * OutThreads->z;
            long toExec = 0;
            int maxThreads = maxBlocks * threads;
            int i, step = 1;
            for (i = loop->dimension - 1; i >= 0; i--) {
                if (OutHighValues[i] != loop->loopBounds[i][1]) {
                    OutLowValues[i] = OutHighValues[i] + loop->loopBounds[i][2];
                    break;
                } else {
                    OutLowValues[i] = loop->loopBounds[i][0];
                }
            }
            for (i = loop->dimension - 1; i >= 0; i--) {
                long rest = (loop->loopBounds[i][1] - OutLowValues[i]) / loop->loopBounds[i][2] + 1;
                long toGet;
                if (step * rest <= maxThreads) {
                    toGet = rest;
                } else if (step * ((rest + 1) / 2) <= maxThreads) {
                    toGet = (rest + 1) / 2;
                } else {
                    toGet = maxThreads / step;
                }
                toExec = step * toGet;
                OutHighValues[i] = OutLowValues[i] + loop->loopBounds[i][2] * (toGet - 1);
                if (OutLowValues[i] != loop->loopBounds[i][0] || OutHighValues[i] != loop->loopBounds[i][1])
                    break;
                step *= (loop->loopBounds[i][1] - loop->loopBounds[i][0]) / loop->loopBounds[i][2] + 1;
            }
            toExec = (toExec + threads - 1) / threads;
            OutBlocks->x = toExec;
            OutBlocks->y = 1;
            OutBlocks->z = 1;
            loop->latestBlocks = toExec;
            loop->restBlocks -= toExec;
            fprintf(stderr, "nsOK7 block=(%d,%d,%d) grid=(%d,%d,%d)\n", OutThreads->x, OutThreads->y, OutThreads->z, OutBlocks->x, OutBlocks->y, OutBlocks->z);
            return 1;
        }
    }
    if (loop->counter == 2) {
        finish_pl(*InDvmGpuLoop);
        fprintf(stderr, "nsOK8\n");
        return 0;
    }
    fprintf(stderr, "nsneOK78\n");
    return 0;
}
*/

static void end_region(DvmGpuRegionRef region) {
    int i;
    for (i = 0; i < region->myArraysCount; i++) {
        DvmGpuArrayHeaderStruct *array = region->myArrays[i];
        fprintf(stderr, "array->devAddr=%ld array->copyout=%d\n", (long)array->deviceAddr, array->copyoutFlag);
        if (array->deviceAddr) {
            if (array->copyoutFlag) {
                perform_copyout(array);
            }
            cudaFree(array->deviceAddr);
            array->deviceAddr = 0;
        }
        safe_free((void **)&array->havePortion);
        safe_free((void **)&array->needPortion);
        fill_headers(array);
        safe_free((void **)&array->additionalHeaders);
        safe_free((void **)&array);
    }
    safe_free((void **)&region->myArrays);
    safe_free((void **)&region->arrays);
    for (i = 0; i < region->additionalHeadersCount; i++) {
        long *header = region->additionalHeaders[i];
        DvmGpuArrayHeaderStruct *array = (DvmGpuArrayHeaderStruct *)header[0];
        int j;
        for (j = 0; j < array->additionalHeadersCount; j++)
            if (array->additionalHeaders[j] == header) {
                array->additionalHeadersCount--;
                if (j < array->additionalHeadersCount) {
                    array->additionalHeaders[j] = array->additionalHeaders[array->additionalHeadersCount];
                }
                j--;
            }
        if (!array->additionalHeadersCount)
            safe_free((void **)&array->additionalHeaders);
    }
    innermostRegion = region->parentRegion;
    safe_free((void **)&region);
}

// Освобождение ненужных структур, подчистка всего, что осталось от вычислительного региона
void end_region_gpu_(DvmGpuComputeRegionRef *InOutComputeRegionGpu) {
    end_region(*InOutComputeRegionGpu);
    (*(long *)InOutComputeRegionGpu) = 0;
}

// Освобождение ненужных структур, подчистка всего, что осталось от региона данных
void end_data_region_gpu_(DvmGpuDataRegionRef *InOutDataRegionGpu) {
    end_region(*InOutDataRegionGpu);
    (*(long *)InOutDataRegionGpu) = 0;
}

/*void set_fixed_block_(char *arg) {
    //x y z
    if (strchr(arg, 'x')) {
        char *s = strchr(arg, 'x');
        *s = 0;
        fixedBlock.x = atoi(arg);
        *s = 'x';
        arg = s + 1;
        s = strchr(arg, 'x');
        *s = 0;
        fixedBlock.y = atoi(arg);
        *s = 'x';
        arg = s + 1;
        fixedBlock.z = atoi(arg);
        arg = s + 1;
    }
}

void set_fixed_block_int_(int *px, int *py, int *pz) {
    fixedBlock.x = *px;
    fixedBlock.y = *py;
    fixedBlock.z = *pz;
}
*/
